  \redish{//blocks shared variable}
  __device__ \greenish{int} \brownish{g_mutex};

  \redish{//centralized barrier function}
  __device__ \greenish{void} \bluish{__gpu_sync}(\greenish{int} \brownish{goalVal})\{
      \greenish{int} \brownish{tid_in_block} = threadIdx.x * blockDim.y + 
         threadIdx.y;

      \redish{// in each block, only thread 0 synchronizes}
      \redish{// with the other blocks}
      \purplish{if} (tid_in_block == 0) \{
         atomicAdd(&g_mutex, 1);
         \redish{// wait for the other blocks}
         \purplish{while}(g_mutex != goalVal) \{ \}
      \}

      \redish{// synchronize within the block}
      __synchthreads();
  \}
